#include"fftz.h"
#include<stdio.h>
#include<string.h>
#include<stdlib.h>
#include<math.h>

int initFftzEnv(FftzEnv* cl){
	int num_platforms;
	int num_devices;
	int err;
	int i, j;
	char* str;
	size_t str_size;
	char device_name_Turks[] = "Turks";
	char device_name_GeForce[] = "GeForce";
	char device_name_Adreno[] = "Adreno";
	
	//get platform ids
	err = clGetPlatformIDs(0, NULL, &num_platforms);
	if(err < 0){
		printf("Error: Calling clGetPlatformIDs the first time failed!\n");
		return err;
	}
	cl->platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
	err = clGetPlatformIDs(num_platforms, cl->platforms, NULL);
	if(err < 0){
		free(cl->platforms);
		printf("Error: Calling clGetPlatformIDs the second time failed!\n");
		return err;
	}
	
	for(i = 0; i < num_platforms; i++){
		err = clGetDeviceIDs(cl->platforms[i], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);
		if(err < 0){
			printf("Error: Calling clGetDeviceIDs the first time failed.\n");
			free(cl->platforms);
			return err;
		}
		cl->devices = (cl_device_id*)malloc(sizeof(cl_device_id) * num_devices);
		clGetDeviceIDs(cl->platforms[i], CL_DEVICE_TYPE_ALL, num_devices, cl->devices, NULL);
		if(err < 0){
			printf("Error: Calling clGetDeviceIDs the second time failed.\n");
			free(cl->devices);
			free(cl->platforms);
			return err;
		}

		for(j = 0; j < num_devices; j++){
			clGetDeviceInfo(cl->devices[j], CL_DEVICE_NAME, 0, NULL, &str_size);
			str = (cl_char*)malloc(sizeof(cl_char) * str_size + 1);
			clGetDeviceInfo(cl->devices[j], CL_DEVICE_NAME, str_size, str, NULL);
			if((strstr(str, device_name_Turks) != NULL) || (strstr(str, device_name_GeForce) != NULL) || (strstr(str, device_name_Adreno) != NULL)){
				printf("Selected device info:\n");
				free(str);

				clGetPlatformInfo(cl->platforms[i], CL_PLATFORM_NAME, 0, NULL, &str_size);
				str = (cl_char*)malloc(sizeof(cl_char) * str_size + 1);
				clGetPlatformInfo(cl->platforms[i], CL_PLATFORM_NAME, str_size, str, NULL);
				printf("  Platform Name: %s\n", str);
				free(str);

				clGetDeviceInfo(cl->devices[j], CL_DEVICE_NAME, 0, NULL, &str_size);
				str = (cl_char*)malloc(sizeof(cl_char) * str_size + 1);
				clGetDeviceInfo(cl->devices[j], CL_DEVICE_NAME, str_size, str, NULL);
				printf("  Device Name: %s\n", str);
				free(str);

				clGetDeviceInfo(cl->devices[j], CL_DEVICE_VENDOR, 0, NULL, &str_size);
				str = (cl_char*)malloc(sizeof(cl_char) * str_size + 1);
				clGetDeviceInfo(cl->devices[j], CL_DEVICE_VENDOR, str_size, str, NULL);
				printf("  Device Vendor: %s\n", str);
				free(str);

				cl->platform = &cl->platforms[i];
				cl->device = &cl->devices[j];

				break;
			}
			free(str);
		}
	}
	
	//create opencl host context
	cl->context = clCreateContext(NULL, 1, cl->device, NULL, NULL, &err);
	if(err < 0){
		free(cl->platforms);
		free(cl->devices);
		printf("Error: Calling clCreateContext failed!\n");
		return err;
	}
	
	//create opencl command queue
	cl->queue = clCreateCommandQueue(cl->context, cl->device[0], CL_QUEUE_PROFILING_ENABLE, &err);
	if(err < 0){
		free(cl->platforms);
		free(cl->devices);
		clReleaseContext(cl->context);
		printf("Error: Calling clCreateCommandQueue failed!\n");
		return err;
	}
	
	//init succeeded
	return 0;
}

int releaseFftzEnv(FftzEnv* cl){
	clReleaseCommandQueue(cl->queue);
	clReleaseContext(cl->context);
	free(cl->devices);
	free(cl->platforms);
	return 0;
}

static int getKernelWorkGroupInfo(FftzEnv* cl, cl_kernel kernel, size_t* kernelWorkGroupSize, size_t* kernelWorkGroupSizeMultiple){
	size_t groupSize;
	int err;
	
	err = clGetKernelWorkGroupInfo(kernel, cl->device[0], CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(groupSize), &groupSize, NULL);
	if(err < 0){
		printf("Failed to get CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE information, function exiting..\n");
		return -1;
	}
	*kernelWorkGroupSizeMultiple = groupSize;
	
	err = clGetKernelWorkGroupInfo(kernel, cl->device[0], CL_KERNEL_WORK_GROUP_SIZE, sizeof(groupSize), &groupSize, NULL);
	if(err < 0){
		printf("Failed to get CL_KERNEL_WORK_GROUP_SIZE information, function exiting..\n");
		return -1;
	}
	*kernelWorkGroupSize = groupSize;
	
	return 0;
}

int defaultGroupSize(int numWorkItems, int maxGroupSize, int multipleGroupSize){
	int groupSize = maxGroupSize;
	
	if(multipleGroupSize <= 0)
		multipleGroupSize = 1;
	while((numWorkItems / maxGroupSize * maxGroupSize != numWorkItems) || (maxGroupSize / multipleGroupSize * multipleGroupSize != maxGroupSize)){
		maxGroupSize--;
		if(maxGroupSize == 0){
			while(numWorkItems / groupSize * groupSize != numWorkItems)
				groupSize--;
			return groupSize;
		}
	}
	return maxGroupSize;
}

int getFftzKernelFromSource(FftzEnv* cl, char* source_str, char* fun_name, FftzKernel* ke){
	cl_program myprogram;
	size_t program_build_log_size;
	char* program_build_log;
	int err;
	
	//compile program
	myprogram = clCreateProgramWithSource(cl->context, 1, (const char**)&source_str, NULL, &err);
	err += clBuildProgram(myprogram, 1, cl->device, NULL, NULL, NULL);
	if(err < 0){
		//print program compiling compiling information
		clGetProgramBuildInfo(myprogram, cl->device[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &program_build_log_size);
		program_build_log = (char*)malloc(sizeof(char) * program_build_log_size);
		clGetProgramBuildInfo(myprogram, cl->device[0], CL_PROGRAM_BUILD_LOG, program_build_log_size, program_build_log, NULL);
		printf("%s\n", program_build_log);
		free(program_build_log);
		
		printf("Error: Compiling program failed!\n");
		return err;
	}
	
	//get kernel
	ke->kernel = NULL;
	ke->kernel = clCreateKernel(myprogram, fun_name, &err);
	if(err < 0){
		printf("Failed to create kernel %s, function exiting...\n", fun_name);
		clReleaseProgram(myprogram);
		return err;
	}
	
	err = getKernelWorkGroupInfo(cl, ke->kernel, &ke->groupSize, &ke->groupSizeMul);
	if(err < 0){
		printf("kernelWorkGroupSize error, function exiting..\n");
		clReleaseKernel(ke->kernel);
		clReleaseProgram(myprogram);
		return err;
	}
	
	clReleaseProgram(myprogram);
	return 0;
}

int releaseFftzKernel(FftzKernel* ke){
	if(ke->kernel != NULL)
		clReleaseKernel(ke->kernel);
	return 0;
}

int releaseFftzFactor(FftzFactor* factor){
	if(factor->data != NULL)
		free(factor->data);
}

int getFftzFactor(FftzFactor* factor, int twiddleOrNot, int length, int totalLength){
	int i, j;
	int temp;

	if(twiddleOrNot == 0){
		factor->dataSize = (sizeof(float) * length)<<2;
		factor->data = NULL;
		factor->data = (float*)malloc(sizeof(float) * length * 4);
		if(factor->data == NULL){
			printf("ERROR: Failed to malloc factor.\n");
			releaseFftzFactor(factor);
			return -1;
		}
		for(i = 0; (1 << i) <= length; i++){
			temp = 1 << i;
			for(j = 0; j < temp; j++){
				factor->data[temp + j] = cos(-FFTZ_PI * 2 * j / temp);
				factor->data[temp + j + length * 2] = sin(-FFTZ_PI * 2 * j / temp);
			}
		}
	}
	else{
		factor->dataSize = (sizeof(float) * length)<<2;
		factor->data = NULL;
		factor->data = (float*)malloc(sizeof(float) * length * 4);
		if(factor->data == NULL){
			printf("ERROR: Failed to malloc factor.\n");
			releaseFftzFactor(factor);
			return -1;
		}
		for(i = 0; i < length; i++){
			factor->data[i] = cos(-FFTZ_PI * 2 * i / totalLength);
			factor->data[i + length * 2] = sin(-FFTZ_PI * 2 * i / totalLength);
			factor->data[i + length] = cos(-FFTZ_PI * 2 * i * length / totalLength);
			factor->data[i + length * 3] = sin(-FFTZ_PI * 2 * i * length / totalLength);
		}
	}
	return 0;
}

int releaseFftzPlan(FftzPlan* plan){
	int i;
	
	if(plan->kernels != NULL){
		for(i = 0; i < plan->numKernels; i++)
			releaseFftzKernel(&plan->kernels[i]);
		free(plan->kernels);
	}
	if(plan->factors != NULL){
		for(i = 0; i < plan->numFactors; i++)
			releaseFftzFactor(&plan->factors[i]);
		free(plan->factors);
	}
	if(plan->steps != NULL)
		free(plan->steps);
	return 0;
}

int initFftzPlan(FftzEnv* cl, FftzPlan* plan, int I){
	int i, j;
	int temp;
	int err;
	
	plan->length = 1 << I;
	plan->lengthLog = I;
	

	//init kernels
	plan->numKernels = 5;
	plan->kernels = NULL;
	plan->kernels = (FftzKernel*)malloc(sizeof(FftzKernel) * plan->numKernels);
	if(plan->kernels == NULL){
		printf("ERROR: Failed to malloc space for plan->kernels.\n");
		releaseFftzPlan(plan);
		return -1;
	}
	
	err = getFftzKernelFromSource(cl, strKernelR2, "gpuFftBase", &plan->kernels[0]);
	if(err < 0){
		printf("ERROR: Failed to create kernel gpuFftBase.\n");
		releaseFftzPlan(plan);
		return err;
	}
	plan->kernels[0].radix = 2;
	plan->kernels[0].groupSizeSelected = defaultGroupSize(plan->length / plan->kernels[0].radix, plan->kernels[0].groupSize, plan->kernels[0].groupSizeMul);

	err = getFftzKernelFromSource(cl, strKernelR4, "gpuFftFour", &plan->kernels[1]);
	if(err < 0){
		printf("Failed to create kernel gpuFftFour, function initFftzPlan exitting..\n");
		releaseFftzPlan(plan);
		return err;
	}
	plan->kernels[1].radix = 4;
	plan->kernels[1].groupSizeSelected = defaultGroupSize(plan->length / plan->kernels[1].radix, plan->kernels[1].groupSize, plan->kernels[1].groupSizeMul);

	err = getFftzKernelFromSource(cl, strKernelR8, "gpuFftOct", &plan->kernels[2]);
	if(err < 0){
		printf("Failed to create kernel gpuFftOct, function initFftzPlan exitting..\n");
		releaseFftzPlan(plan);
		return err;
	}
	plan->kernels[2].radix = 8;
	plan->kernels[2].groupSizeSelected = defaultGroupSize(plan->length / plan->kernels[2].radix, plan->kernels[2].groupSize, plan->kernels[2].groupSizeMul);

	err = getFftzKernelFromSource(cl, strKernelTwiddle, "gpuFftTwiddle", &plan->kernels[3]);
	if(err < 0){
		printf("Failed to create kernel gpuFftTwiddle, function initFftzPlan exitting...\n");
		releaseFftzPlan(plan);
		return err;
	}
	plan->kernels[3].radix = 1;
	plan->kernels[3].groupSizeSelected = defaultGroupSize(plan->length, plan->kernels[3].groupSize, plan->kernels[3].groupSizeMul);

	err = getFftzKernelFromSource(cl, strKernelPermutation, "gpuFftPermutation", &plan->kernels[4]);
	if(err < 0){
		printf("Failed to create kernel gpuFftPermutation, function initFftzPlan exitting...\n");
		releaseFftzPlan(plan);
		return err;
	}
	plan->kernels[4].radix = 1;
	plan->kernels[4].groupSizeSelected = defaultGroupSize(plan->length, plan->kernels[4].groupSize, plan->kernels[4].groupSizeMul);
	

	//do not need the large single-vector algrithm
	if(plan->length <= NO_TWIDDLE_LENGTH){ 
		//get number of steps
		plan->numSteps = 0;
		i = I;
		while(i > 0){
			if(i >= 2){
				i -= 2;
			}
			else{
				i -= 1;
			}
			plan->numSteps++;
		}
		//get factors
		plan->numFactors = 1;
		plan->factors = NULL;
		plan->factors = (FftzFactor*)malloc(sizeof(FftzFactor) * plan->numFactors);
		if(plan->factors == NULL){
			printf("ERROR: Failed to malloc plan->factors.\n");
			releaseFftzPlan(plan);
			return -1;
		}
		getFftzFactor(&plan->factors[0], 0, plan->length, 1);
		plan->maxFactorDataSize = plan->factors[0].dataSize;
		//get steps
		plan->steps = NULL;
		plan->steps = (FftzStep*)malloc(sizeof(FftzStep) * plan->numSteps);
		if(plan->steps == NULL){
			printf("ERROR: Failed to malloc plan->steps.\n");
			releaseFftzPlan(plan);
			return -1;
		}
		j = 0;
		for(i = 0; i < plan->numSteps; i++){
			plan->steps[i].lengthLog = j;
			plan->steps[i].length = (1 << j);
			plan->steps[i].partTotalLengthLog = j;
			plan->steps[i].partTotalLength = (1 << j);
			plan->steps[i].type = 0;
			plan->steps[i].factor = &plan->factors[0];
			if(I - j >= 2){
				plan->steps[i].kernel = &plan->kernels[1];
				j += 2;
			}
			else{
				plan->steps[i].kernel = &plan->kernels[0];
				j += 1;
			}
		}		
	}
	//a large single-vector algrithm is needed
	else{
		int lengthLeft, lengthRight;
		int numStepsLeft, numStepsRight;
		int lengthLeftLogBy2, lengthRightLogBy2;

		//get numStepsLeft and numStepsRight
		i = 0;
		numStepsLeft = 0;
		numStepsRight = 0;
		lengthLeftLogBy2 = 0;
		lengthRightLogBy2 = 0;
		while(1){
			if((I - lengthLeftLogBy2 - lengthRightLogBy2) > 4){
				lengthLeftLogBy2 += 2;
				lengthRightLogBy2 += 2;
				numStepsLeft ++;
				numStepsRight ++;
				continue;
			}
			else{
				switch(I - lengthLeftLogBy2 - lengthRightLogBy2){
					case 4:
						lengthLeftLogBy2 += 2;
						lengthRightLogBy2 += 2;
						numStepsLeft ++;
						numStepsRight ++;
						break;
					case 3:
						lengthLeftLogBy2 += 2;
						lengthRightLogBy2 += 1;
						numStepsLeft ++;
						numStepsRight ++;
						break;
					case 2:
						lengthLeftLogBy2 += 2;
						numStepsLeft ++;
						break;
					case 1:
						lengthLeftLogBy2 += 1;
						numStepsLeft ++;
						break;
					default:
						break;
				}
				break;
			}
		}
		printf("Left-Right info:\n");
		printf("  lengthLeftLogBy2: %d\n", lengthLeftLogBy2);
		printf("  lengthRightLogBy2: %d\n", lengthRightLogBy2);
		printf("  numStepsLeft: %d\n", numStepsLeft);
		printf("  numStepsRight: %d\n", numStepsRight);
		//get factors
		plan->numFactors = 3;
		plan->factors = NULL;
		plan->factors = (FftzFactor*)malloc(sizeof(FftzFactor) * plan->numFactors);
		if(plan->factors == NULL){
			printf("ERROR: Failed to malloc plan->factors.\n");
			releaseFftzPlan(plan);
			return -1;
		}
		getFftzFactor(&plan->factors[0], 0, (1<<lengthLeftLogBy2), 1);
		getFftzFactor(&plan->factors[1], 1, (1<<lengthLeftLogBy2), plan->length);
		getFftzFactor(&plan->factors[2], 0, (1<<lengthRightLogBy2), 1);
		plan->maxFactorDataSize = 0;
		for(i = 0; i < plan->numFactors; i++)
			if(plan->maxFactorDataSize < plan->factors[i].dataSize)
				plan->maxFactorDataSize = plan->factors[i].dataSize;
		//get steps
		plan->numSteps = numStepsLeft + numStepsRight + 2;
		plan->steps = NULL;
		plan->steps = (FftzStep*)malloc(sizeof(FftzStep) * plan->numSteps);
		if(plan->steps == NULL){
			printf("ERROR: Failed to malloc plan->steps.\n");
			releaseFftzPlan(plan);
			return -1;
		}
		j = 0;
		for(i = 0; i < numStepsLeft; i++){
			plan->steps[i].lengthLog = j;
			plan->steps[i].length = (1 << j);
			plan->steps[i].partTotalLengthLog = lengthLeftLogBy2;
			plan->steps[i].partTotalLength = (1 << lengthLeftLogBy2);
			plan->steps[i].type = 0;
			plan->steps[i].factor = &plan->factors[0];
			if(lengthLeftLogBy2 - j >= 2){
				plan->steps[i].kernel = &plan->kernels[1];
				j += 2;
			}
			else{
				plan->steps[i].kernel = &plan->kernels[0];
				j += 1;
			}
		}
		plan->steps[numStepsLeft].lengthLog = j;
		plan->steps[numStepsLeft].length = (1<<j);
		plan->steps[numStepsLeft].type = 1;
		plan->steps[numStepsLeft].kernel = &plan->kernels[3];
		plan->steps[numStepsLeft].factor = &plan->factors[1];
		j = 0;
		for(i = 0; i < numStepsRight; i++){
			plan->steps[i + numStepsLeft + 1].lengthLog = j;
			plan->steps[i + numStepsLeft + 1].length = (1 << j);
			plan->steps[i + numStepsLeft + 1].partTotalLengthLog = lengthRightLogBy2;
			plan->steps[i + numStepsLeft + 1].partTotalLength = (1 << lengthRightLogBy2);
			plan->steps[i + numStepsLeft + 1].type = 0;
			plan->steps[i + numStepsLeft + 1].factor = &plan->factors[2];
			if(lengthRightLogBy2 - j >= 2){
				plan->steps[i + numStepsLeft + 1].kernel = &plan->kernels[1];
				j += 2;
			}
			else{
				plan->steps[i + numStepsLeft + 1].kernel = &plan->kernels[0];
				j += 1;
			}
		}
		plan->steps[numStepsLeft + numStepsRight + 1].lengthLog = lengthRightLogBy2;
		plan->steps[numStepsLeft + numStepsRight + 1].length = (1<<lengthRightLogBy2);
		plan->steps[numStepsLeft + numStepsRight + 1].type = 2;
		plan->steps[numStepsLeft + numStepsRight + 1].kernel = &plan->kernels[4];
		plan->steps[numStepsLeft + numStepsRight + 1].factor = NULL;
	}

	return 0;
}



int fftzRun(FftzEnv* cl, FftzPlan* plan, float* data0){
	cl_mem buffer_data0 = NULL;
	cl_mem buffer_data1 = NULL;
	cl_mem buffer_factors = NULL;
	FftzFactor* lastFactor = NULL;
	FftzStep* step = NULL;
	int i;
	int even_odd = 0;
	int Ns;
	size_t globalws[1];
	size_t localws[1];
	cl_event myevent;
	cl_ulong task_queued, task_start, task_end;
	double total_time = 0;
	int err;
	
	//create buffers
	buffer_data0 = clCreateBuffer(cl->context, CL_MEM_READ_WRITE, sizeof(float) * plan->length * 2, NULL, &err);
	if(err < 0){
		printf("Failed to create buffer_data0 with code %d.\n", err);
		return err;
	}
	buffer_data1 = clCreateBuffer(cl->context, CL_MEM_READ_WRITE, sizeof(float) * plan->length * 2, NULL, &err);
	if(err < 0){
		printf("Failed to create buffer_data1 with code %d.\n", err);
		clReleaseMemObject(buffer_data0);
		return err;
	}
	buffer_factors = clCreateBuffer(cl->context, CL_MEM_READ_ONLY, plan->maxFactorDataSize, NULL, &err);
	if(err < 0){
		printf("Failed to create buffer_factors with error code %d.\n", err);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		return err;
	}

	//write buffers data0
	err = clEnqueueWriteBuffer(cl->queue, buffer_data0, CL_FALSE, 0, sizeof(float) * plan->length * 2, data0, 0, NULL, NULL);
	if(err < 0){
		printf("Failed to write data to buffer_data0 with code %d.\n", err);
		clReleaseMemObject(buffer_factors);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		return err;
	}
		
	for(i = 0; i < plan->numSteps; i++){
		step = &plan->steps[i];
		if((lastFactor != step->factor) && (step->factor != NULL)){
			err = clEnqueueWriteBuffer(cl->queue, buffer_factors, CL_FALSE, 0, step->factor->dataSize, step->factor->data, 0, NULL, NULL);
			if(err < 0){
				printf("Failed to write data to buffer_factors with code %d.\n", err);
				clReleaseMemObject(buffer_factors);
				clReleaseMemObject(buffer_data1);
				clReleaseMemObject(buffer_data0);
				return err;
			}
			lastFactor = step->factor;
		}

		switch(step->type){
			case 0:
				err += clSetKernelArg(step->kernel->kernel, even_odd, sizeof(cl_mem), &buffer_data0);
				err += clSetKernelArg(step->kernel->kernel, 1 - even_odd, sizeof(cl_mem), &buffer_data1);
				err += clSetKernelArg(step->kernel->kernel, 2, sizeof(cl_mem), &buffer_factors);
				err += clSetKernelArg(step->kernel->kernel, 3, sizeof(int), &plan->length);
				err += clSetKernelArg(step->kernel->kernel, 4, sizeof(int), &step->length);
				err += clSetKernelArg(step->kernel->kernel, 5, sizeof(int), &step->partTotalLength);
				even_odd = 1 - even_odd;
				break;
			case 1:
				if(even_odd == 0)
					err += clSetKernelArg(step->kernel->kernel, 0, sizeof(cl_mem), &buffer_data0);
				else
					err += clSetKernelArg(step->kernel->kernel, 0, sizeof(cl_mem), &buffer_data1);
				err += clSetKernelArg(step->kernel->kernel, 1, sizeof(cl_mem), &buffer_factors);
				err += clSetKernelArg(step->kernel->kernel, 2, sizeof(int), &plan->length);
				err += clSetKernelArg(step->kernel->kernel, 3, sizeof(int), &step->length);
				break;
			case 2:
				Ns = plan->length / step->length;
				err += clSetKernelArg(step->kernel->kernel, even_odd, sizeof(cl_mem), &buffer_data0);
				err += clSetKernelArg(step->kernel->kernel, 1 - even_odd, sizeof(cl_mem), &buffer_data1);
				err += clSetKernelArg(step->kernel->kernel, 2, sizeof(int), &plan->length);
				err += clSetKernelArg(step->kernel->kernel, 3, sizeof(int), &Ns);
				err += clSetKernelArg(step->kernel->kernel, 4, sizeof(int), &step->length);
				even_odd = 1 - even_odd;
				break;
			default:
				printf("Unknown step type when setting kernel args.\n");
				break;
		}

		if(err < 0){
			printf("Failed to set args of kernel with code %d.\n", err);
			clReleaseMemObject(buffer_factors);
			clReleaseMemObject(buffer_data1);
			clReleaseMemObject(buffer_data0);
			return err;
		}

		globalws[0] = plan->length / step->kernel->radix;
		localws[0] = step->kernel->groupSizeSelected;
		err = clEnqueueNDRangeKernel(cl->queue, step->kernel->kernel, 1, NULL, globalws, localws, 0, NULL, &myevent);
		if(err < 0){
			printf("Failed to enqueue kernel  with code %d.\n", err);
			clReleaseMemObject(buffer_factors);
			clReleaseMemObject(buffer_data1);
			clReleaseMemObject(buffer_data0);
			return err;
		}
		clFinish(cl->queue);
		clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &task_start, NULL);
		clGetEventProfilingInfo(myevent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &task_end, NULL);
		total_time += (double)(task_end - task_start) / 1000.0;	//us
		clReleaseEvent(myevent);
	}

	if(even_odd == 1)
		err = clEnqueueReadBuffer(cl->queue, buffer_data1, CL_FALSE, 0, sizeof(float) * plan->length * 2, data0, 0, NULL, NULL);
	else
		err = clEnqueueReadBuffer(cl->queue, buffer_data0, CL_FALSE, 0, sizeof(float) * plan->length * 2, data0, 0, NULL, NULL);
	if(err < 0){
		printf("Failed to fetch results with error code %d.\n", err);
		clReleaseMemObject(buffer_factors);
		clReleaseMemObject(buffer_data1);
		clReleaseMemObject(buffer_data0);
		return err;
	}
	clFinish(cl->queue);


	clReleaseMemObject(buffer_factors);
	clReleaseMemObject(buffer_data1);
	clReleaseMemObject(buffer_data0);
	return (int)total_time;
}

static int mystrncpy(char *dest, char *src, int n){
	int i;

	for(i = 0; i < n; i++)
		*dest++ = *src++;
	*dest = '\0';
	return 0;
}

int strrep(char *src, char * src_temp, char *match, char *replace){
	char *find_pos;

	find_pos = strstr(src, match);
	if((!find_pos) || (!match))
		return -1;
	while(find_pos){
		mystrncpy(src_temp, src, find_pos - src);
		strcat(src_temp, replace);
		strcat(src_temp, find_pos + strlen(match));
		strcpy(src, src_temp);

		find_pos = strstr(src, match);
	}

	return 0;
}
